/*
目的:
    大量数据 并行计算结果
*/

#include <stdio.h>
#include <string>
#include <time.h>
#include <sys/time.h>
#include <math.h>
#include <stdlib.h>
#include <cuda_runtime.h>
using namespace std;

#define GPU_TASK_THREADS 256
float gpu_clockRate;

// cuda核函数: 单线程
__global__ void mem_serial_read(float* array, int size, double* result, clock_t* clock_consum)
{
    clock_t clock_start = clock();

    double sum = 0.0;
    for (int i = 0; i < size; i++) {
        sum += sqrtf(array[i]);
        sum += (array[i] * array[i] * array[i]);
    }
    result[0] = sum;
    *clock_consum = clock() - clock_start;
}

// cuda核函数: 多线程, 线程按内存顺序读取
__global__ void mem_sequential_read(float* array, int size, double* result, clock_t* clock_consum)
{
    // cuda线程id
    const int tid = threadIdx.x;
    clock_t clock_start;
    if (tid == 0) {
        clock_start = clock();
    }

    // 要计算的元素个数
    const int cal_size = size / GPU_TASK_THREADS;
    double sum = 0.0;

    for (int i = tid * cal_size; i < (tid + 1) * cal_size; i++) {
        sum += sqrtf(array[i]);
        sum += (array[i] * array[i] * array[i]);
    }
    result[tid] = sum;

    if (tid == 0) {
        *clock_consum = clock() - clock_start;
    }
}

// cuda核函数: 多线程, 内存并行顺序读取
__global__ void mem_parallel_read(float* array, int size, double* result, clock_t* clock_consum)
{
    // cuda线程id
    const int tid = threadIdx.x;
    clock_t clock_start;
    if (tid == 0) {
        clock_start = clock();
    }

    // 要计算的元素个数
    const int cal_size = size / GPU_TASK_THREADS;
    double sum = 0.0;

    for (int i = tid; i < size; i += GPU_TASK_THREADS) {
        sum += sqrtf(array[i]);
        sum += (array[i] * array[i] * array[i]);
    }
    result[tid] = sum;

    if (tid == 0) {
        *clock_consum = clock() - clock_start;
    }
}

float gpu_parallel_test(float* array, int size, string type = "")
{
    // host 变量
    double result[GPU_TASK_THREADS];
    clock_t clock_consum = 0;

    // device 变量
    float* d_array;
    double* d_result;
    clock_t* d_clock_consum;

    // cauda device申请内存
    cudaMalloc((void**)&d_array, size * sizeof(float));
    cudaMalloc((void**)&d_result, GPU_TASK_THREADS * sizeof(double));
    cudaMalloc((void**)&d_clock_consum, sizeof(clock_t));

    // 将host数据拷贝到device
    cudaMemcpy(d_array, array, size * sizeof(float), cudaMemcpyHostToDevice);

    // 调用核函数
    if (type == "" || type == "signal_thread") {
        mem_serial_read << <1, 1 >> > (d_array, size, d_result, d_clock_consum);
    }
    else if (type == "mem_sequential_read") {
        mem_sequential_read << <1, GPU_TASK_THREADS, 0 >> > (d_array, size, d_result, d_clock_consum);
    }
    else if (type == "mem_parallel_read") {
        mem_parallel_read << <1, GPU_TASK_THREADS, 0 >> > (d_array, size, d_result, d_clock_consum);
    }

    // 将device数据拷贝到host
    cudaMemcpy(&result, d_result, GPU_TASK_THREADS * sizeof(double), cudaMemcpyDeviceToHost);
    cudaMemcpy(&clock_consum, d_clock_consum, sizeof(clock_t), cudaMemcpyDeviceToHost);


    float sum = 0.0;
    for (int i = 0; i < GPU_TASK_THREADS; i++) {
        sum += result[i];
    }
    printf("gpu %20s:%f, 耗时:%f\n", type.c_str(), sum, clock_consum / gpu_clockRate);

    // 释放device内存
    cudaFree(d_array);
    cudaFree(d_result);
    cudaFree(d_clock_consum);

    return sum;
}


float cpu_test(float* array, int size)
{
    struct timeval start_tv, end_tv, consum_tv;
    gettimeofday(&start_tv, NULL);
    double sum = 0.0;
    for (int i = 0; i < size; i++) {
        sum += sqrtf(array[i]);
        sum += (array[i] * array[i] * array[i]);
    }

    gettimeofday(&end_tv, NULL);
    timersub(&end_tv, &start_tv, &consum_tv);
    double consum_time = consum_tv.tv_sec + consum_tv.tv_usec / 1000000.0;

    printf("cpu 计算结果%12s:%f, 耗时:%lf\n", "", sum,consum_time);
    return sum;
}

void init_gpu_clockRate()
{
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    gpu_clockRate = prop.clockRate * 1000.0;
}

float* gen_random_data(int size)
{
    float* data = new float[size];
    for (int i = 0; i < size; i++) {
        data[i] = rand() % 1000;
    }

    return data;
}

int main()
{
    init_gpu_clockRate();

    // 测试数据
    int size = 1024 * 1024;
    float* array = gen_random_data(size);

    gpu_parallel_test(array, size, "signal_thread");
    gpu_parallel_test(array, size, "mem_sequential_read");
    gpu_parallel_test(array, size, "mem_parallel_read");
    cpu_test(array, size);

    // 释放内存
    delete[]array;

    return 0;
}

/* 运行结果:
gpu        signal_thread:261707205181440.000000, 耗时:0.133453
gpu  mem_sequential_read:261707205181440.000000, 耗时:0.001345
gpu    mem_parallel_read:261707121295360.000000, 耗时:0.001407
cpu 计算结果            :261707198325593.750000, 耗时:0.003926

 */